<!-- HTML header for doxygen 1.8.3.1-->
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<meta name="generator" content="Doxygen 1.8.4"/>
<title>CUB: cub::WarpReduce&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt; Class Template Reference</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
  $(document).ready(function() { searchBox.OnSelectItem(0); });
</script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
<link href="extra_stylesheet.css" rel="stylesheet" type="text/css"/>
<link rel="shortcut icon" href="favicon.ico" type="image/x-icon" />
<script type="text/javascript">
  var _gaq = _gaq || [];
  _gaq.push(['_setAccount', 'UA-38890655-1']);
  _gaq.push(['_trackPageview']);
  (function() {
    var ga = document.createElement('script'); ga.type = 'text/javascript'; ga.async = true;
    ga.src = ('https:' == document.location.protocol ? 'https://ssl' : 'http://www') + '.google-analytics.com/ga.js';
    var s = document.getElementsByTagName('script')[0]; s.parentNode.insertBefore(ga, s);
  })();
</script>
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
 <tbody>
 <tr style="height: 56px;">
  <td style="padding-left: 0.5em;">
   <div id="projectname">CUB
   </div>
  </td>
 </tr>
 </tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.4 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
  <div id="navrow1" class="tabs">
    <ul class="tablist">
      <li><a href="index.html"><span>Main&#160;Page</span></a></li>
      <li><a href="modules.html"><span>Modules</span></a></li>
      <li class="current"><a href="annotated.html"><span>Classes</span></a></li>
      <li>
        <div id="MSearchBox" class="MSearchBoxInactive">
        <span class="left">
          <img id="MSearchSelect" src="search/mag_sel.png"
               onmouseover="return searchBox.OnSearchSelectShow()"
               onmouseout="return searchBox.OnSearchSelectHide()"
               alt=""/>
          <input type="text" id="MSearchField" value="Search" accesskey="S"
               onfocus="searchBox.OnSearchFieldFocus(true)" 
               onblur="searchBox.OnSearchFieldFocus(false)" 
               onkeyup="searchBox.OnSearchFieldChange(event)"/>
          </span><span class="right">
            <a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
          </span>
        </div>
      </li>
    </ul>
  </div>
  <div id="navrow2" class="tabs2">
    <ul class="tablist">
      <li><a href="annotated.html"><span>Class&#160;List</span></a></li>
      <li><a href="classes.html"><span>Class&#160;Index</span></a></li>
    </ul>
  </div>
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
     onmouseover="return searchBox.OnSearchSelectShow()"
     onmouseout="return searchBox.OnSearchSelectHide()"
     onkeydown="return searchBox.OnSearchSelectKey(event)">
<a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(0)"><span class="SelectionMark">&#160;</span>All</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(1)"><span class="SelectionMark">&#160;</span>Classes</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(2)"><span class="SelectionMark">&#160;</span>Namespaces</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(3)"><span class="SelectionMark">&#160;</span>Files</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(4)"><span class="SelectionMark">&#160;</span>Functions</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(5)"><span class="SelectionMark">&#160;</span>Variables</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(6)"><span class="SelectionMark">&#160;</span>Typedefs</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(7)"><span class="SelectionMark">&#160;</span>Enumerations</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(8)"><span class="SelectionMark">&#160;</span>Enumerator</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(9)"><span class="SelectionMark">&#160;</span>Groups</a></div>

<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0" 
        name="MSearchResults" id="MSearchResults">
</iframe>
</div>

<div id="nav-path" class="navpath">
  <ul>
<li class="navelem"><a class="el" href="namespacecub.html">cub</a></li><li class="navelem"><a class="el" href="classcub_1_1_warp_reduce.html">WarpReduce</a></li>  </ul>
</div>
</div><!-- top -->
<div class="header">
  <div class="summary">
<a href="#nested-classes">Classes</a> &#124;
<a href="classcub_1_1_warp_reduce-members.html">List of all members</a>  </div>
  <div class="headertitle">
<div class="title">cub::WarpReduce&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt; Class Template Reference<div class="ingroups"><a class="el" href="group___warp_module.html">Warp-wide</a></div></div>  </div>
</div><!--header-->
<div class="contents">
<a name="details" id="details"></a><h2 class="groupheader">Detailed description</h2>
<div class="textblock"><h3>template&lt;
    typename T, 
    int LOGICAL_WARPS = 1, 
    int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt;<br/>
class cub::WarpReduce&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;</h3>

<p>The <a class="el" href="classcub_1_1_warp_reduce.html" title="The WarpReduce class provides collective methods for computing a parallel reduction of items partitio...">WarpReduce</a> class provides <a href="index.html#sec0"><em>collective</em></a> methods for computing a parallel reduction of items partitioned across CUDA warp threads. </p>
<div class="image">
<img src="warp_reduce_logo.png" alt="warp_reduce_logo.png"/>
<div class="caption">
.</div></div>
 <dl class="section user"><dt>Overview</dt><dd>A <a href="http://en.wikipedia.org/wiki/Reduce_(higher-order_function)"><em>reduction</em></a> (or <em>fold</em>) uses a binary combining operator to compute a single aggregate from a list of input elements.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">T</td><td>The reduction input/output element type </td></tr>
    <tr><td class="paramname">LOGICAL_WARPS</td><td><b>[optional]</b> The number of entrant "logical" warps performing concurrent warp reductions. Default is 1. </td></tr>
    <tr><td class="paramname">LOGICAL_WARP_THREADS</td><td><b>[optional]</b> The number of threads per "logical" warp (may be less than the number of hardware warp threads). Default is the warp size of the targeted CUDA compute-capability (e.g., 32 threads for SM20).</td></tr>
  </table>
  </dd>
</dl>
<dl class="section user"><dt>Simple Examples</dt><dd>Every thread in the warp uses the <a class="el" href="classcub_1_1_warp_reduce.html" title="The WarpReduce class provides collective methods for computing a parallel reduction of items partitio...">WarpReduce</a> class by first specializing the <a class="el" href="classcub_1_1_warp_reduce.html" title="The WarpReduce class provides collective methods for computing a parallel reduction of items partitio...">WarpReduce</a> type, then instantiating an instance with parameters for communication, and finally invoking collective member functions. </dd></dl>
<dl class="section user"><dt></dt><dd>The code snippet below illustrates four concurrent warp sum reductions within a block of 128 threads (one per each of the 32-thread warps). </dd></dl>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpReduce for 4 warps on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_reduce.html" title="The WarpReduce class provides collective methods for computing a parallel reduction of items partitio...">cub::WarpReduce&lt;int, 4&gt;</a> <a class="code" href="classcub_1_1_warp_reduce.html#aed6c8ead3a4c2fa56ae1d5851fb36848" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpReduce</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpReduce</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpReduce::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain one input item per thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data = ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Return the warp-wide sums to each lane0 (threads 0, 32, 64, and 96)</span></div>
<div class="line">    <span class="keywordtype">int</span> aggregate = <a class="code" href="classcub_1_1_warp_reduce.html#aed6c8ead3a4c2fa56ae1d5851fb36848" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpReduce</a>(temp_storage).Sum(thread_data);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>0, 1, 2, 3, ..., 127</code>. The corresponding output <code>aggregate</code> in threads 0, 32, 64, and 96 will <code>496</code>, <code>1520</code>, <code>2544</code>, and <code>3568</code>, respectively (and is undefined in other threads).</dd></dl>
<dl class="section user"><dt></dt><dd>The code snippet below illustrates a single warp sum reduction within a block of 128 threads. </dd></dl>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpReduce for one warp on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_reduce.html" title="The WarpReduce class provides collective methods for computing a parallel reduction of items partitio...">cub::WarpReduce&lt;int, 1&gt;</a> <a class="code" href="classcub_1_1_warp_reduce.html#aed6c8ead3a4c2fa56ae1d5851fb36848" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpReduce</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpReduce</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpReduce::TempStorage temp_storage;</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Only the first warp performs a reduction</span></div>
<div class="line">    <span class="keywordflow">if</span> (threadIdx.x &lt; 32)</div>
<div class="line">    {</div>
<div class="line">        <span class="comment">// Obtain one input item per thread</span></div>
<div class="line">        <span class="keywordtype">int</span> thread_data = ...</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Return the warp-wide sum to lane0</span></div>
<div class="line">        <span class="keywordtype">int</span> aggregate = <a class="code" href="classcub_1_1_warp_reduce.html#aed6c8ead3a4c2fa56ae1d5851fb36848" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpReduce</a>(temp_storage).Sum(thread_data);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the warp of threads is <code>0, 1, 2, 3, ..., 31</code>. The corresponding output <code>aggregate</code> in thread0 will be <code>496</code> (and is undefined in other threads).</dd></dl>
<dl class="section user"><dt>Usage and Performance Considerations</dt><dd><ul>
<li>Supports "logical" warps smaller than the physical warp size (e.g., logical warps of 8 threads)</li>
<li>The number of entrant threads must be an multiple of <code>LOGICAL_WARP_THREADS</code> </li>
<li>Warp reductions are concurrent if more than one logical warp is participating</li>
<li>Uses special instructions when applicable (e.g., warp <code>SHFL</code> instructions)</li>
<li>Uses synchronization-free communication between warp lanes when applicable</li>
<li>Zero bank conflicts for most types</li>
<li>Computation is slightly more efficient (i.e., having lower instruction overhead) for:<ul>
<li>Summation (<b><em>vs.</em></b> generic reduction)</li>
<li>The architecture's warp size is a whole multiple of <code>LOGICAL_WARP_THREADS</code> </li>
</ul>
</li>
</ul>
</dd></dl>

<p>Definition at line <a class="el" href="warp__reduce_8cuh_source.html#l00141">141</a> of file <a class="el" href="warp__reduce_8cuh_source.html">warp_reduce.cuh</a>.</p>
</div><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="nested-classes"></a>
Classes</h2></td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_warp_reduce_1_1_temp_storage.html">TempStorage</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">The operations exposed by <a class="el" href="classcub_1_1_warp_reduce.html" title="The WarpReduce class provides collective methods for computing a parallel reduction of items partitio...">WarpReduce</a> require a temporary memory allocation of this nested type for thread communication. This opaque storage can be allocated directly using the <code>__shared__</code> keyword. Alternatively, it can be aliased to externally allocated memory (shared or global) or <code>union</code>'d with other storage allocation types to facilitate memory reuse.  <a href="structcub_1_1_warp_reduce_1_1_temp_storage.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="pub-methods"></a>
Public Methods</h2></td></tr>
<tr><td colspan="2"><div class="groupHeader">Collective constructors</div></td></tr>
<tr class="memitem:aed6c8ead3a4c2fa56ae1d5851fb36848"><td class="memItemLeft" align="right" valign="top"><a class="anchor" id="aed6c8ead3a4c2fa56ae1d5851fb36848"></a>
__device__ __forceinline__&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_reduce.html#aed6c8ead3a4c2fa56ae1d5851fb36848">WarpReduce</a> ()</td></tr>
<tr class="memdesc:aed6c8ead3a4c2fa56ae1d5851fb36848"><td class="mdescLeft">&#160;</td><td class="mdescRight">Collective constructor for 1D thread blocks using a private static allocation of shared memory as temporary storage. Logical warp and lane identifiers are constructed from <code>threadIdx.x</code>. <br/></td></tr>
<tr class="separator:aed6c8ead3a4c2fa56ae1d5851fb36848"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a7e026682e53d5af519912ab9a9215e3b"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_reduce.html#a7e026682e53d5af519912ab9a9215e3b">WarpReduce</a> (<a class="el" href="structcub_1_1_warp_reduce_1_1_temp_storage.html">TempStorage</a> &amp;temp_storage)</td></tr>
<tr class="memdesc:a7e026682e53d5af519912ab9a9215e3b"><td class="mdescLeft">&#160;</td><td class="mdescRight">Collective constructor for 1D thread blocks using the specified memory allocation as temporary storage. Logical warp and lane identifiers are constructed from <code>threadIdx.x</code>.  <a href="#a7e026682e53d5af519912ab9a9215e3b">More...</a><br/></td></tr>
<tr class="separator:a7e026682e53d5af519912ab9a9215e3b"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:aef629abf014a0f1f9646d3b9123fcaa6"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_reduce.html#aef629abf014a0f1f9646d3b9123fcaa6">WarpReduce</a> (int warp_id, int lane_id)</td></tr>
<tr class="memdesc:aef629abf014a0f1f9646d3b9123fcaa6"><td class="mdescLeft">&#160;</td><td class="mdescRight">Collective constructor using a private static allocation of shared memory as temporary storage. Threads are identified using the given warp and lane identifiers.  <a href="#aef629abf014a0f1f9646d3b9123fcaa6">More...</a><br/></td></tr>
<tr class="separator:aef629abf014a0f1f9646d3b9123fcaa6"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:aac44f5e8003911f074156366969947e4"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_reduce.html#aac44f5e8003911f074156366969947e4">WarpReduce</a> (<a class="el" href="structcub_1_1_warp_reduce_1_1_temp_storage.html">TempStorage</a> &amp;temp_storage, int warp_id, int lane_id)</td></tr>
<tr class="memdesc:aac44f5e8003911f074156366969947e4"><td class="mdescLeft">&#160;</td><td class="mdescRight">Collective constructor using the specified memory allocation as temporary storage. Threads are identified using the given warp and lane identifiers.  <a href="#aac44f5e8003911f074156366969947e4">More...</a><br/></td></tr>
<tr class="separator:aac44f5e8003911f074156366969947e4"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Summation reductions</div></td></tr>
<tr class="memitem:a3327a2b8c9bdce58f7af1e2485489eed"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__ T&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_reduce.html#a3327a2b8c9bdce58f7af1e2485489eed">Sum</a> (T input)</td></tr>
<tr class="memdesc:a3327a2b8c9bdce58f7af1e2485489eed"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes a warp-wide sum in each active warp. The output is valid in warp <em>lane</em><sub>0</sub>.  <a href="#a3327a2b8c9bdce58f7af1e2485489eed">More...</a><br/></td></tr>
<tr class="separator:a3327a2b8c9bdce58f7af1e2485489eed"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a2ccc6eb62a1e6caf145eb94a8cd58f0f"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__ T&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_reduce.html#a2ccc6eb62a1e6caf145eb94a8cd58f0f">Sum</a> (T input, int valid_items)</td></tr>
<tr class="memdesc:a2ccc6eb62a1e6caf145eb94a8cd58f0f"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes a partially-full warp-wide sum in each active warp. The output is valid in warp <em>lane</em><sub>0</sub>.  <a href="#a2ccc6eb62a1e6caf145eb94a8cd58f0f">More...</a><br/></td></tr>
<tr class="separator:a2ccc6eb62a1e6caf145eb94a8cd58f0f"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:abc0cd1853c09b27fc45f564596163be4"><td class="memTemplParams" colspan="2">template&lt;typename Flag &gt; </td></tr>
<tr class="memitem:abc0cd1853c09b27fc45f564596163be4"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ T&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_reduce.html#abc0cd1853c09b27fc45f564596163be4">HeadSegmentedSum</a> (T input, Flag head_flag)</td></tr>
<tr class="memdesc:abc0cd1853c09b27fc45f564596163be4"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes a segmented sum in each active warp where segments are defined by head-flags. The sum of each segment is returned to the first lane in that segment (which always includes <em>lane</em><sub>0</sub>).  <a href="#abc0cd1853c09b27fc45f564596163be4">More...</a><br/></td></tr>
<tr class="separator:abc0cd1853c09b27fc45f564596163be4"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a5b0134131d03909e43d24d6b0b50beb4"><td class="memTemplParams" colspan="2">template&lt;typename Flag &gt; </td></tr>
<tr class="memitem:a5b0134131d03909e43d24d6b0b50beb4"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ T&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_reduce.html#a5b0134131d03909e43d24d6b0b50beb4">TailSegmentedSum</a> (T input, Flag tail_flag)</td></tr>
<tr class="memdesc:a5b0134131d03909e43d24d6b0b50beb4"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes a segmented sum in each active warp where segments are defined by tail-flags. The sum of each segment is returned to the first lane in that segment (which always includes <em>lane</em><sub>0</sub>).  <a href="#a5b0134131d03909e43d24d6b0b50beb4">More...</a><br/></td></tr>
<tr class="separator:a5b0134131d03909e43d24d6b0b50beb4"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Generic reductions</div></td></tr>
<tr class="memitem:a1aa0a6e1a2c3feabbc20c27864df15c1"><td class="memTemplParams" colspan="2">template&lt;typename ReductionOp &gt; </td></tr>
<tr class="memitem:a1aa0a6e1a2c3feabbc20c27864df15c1"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ T&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_reduce.html#a1aa0a6e1a2c3feabbc20c27864df15c1">Reduce</a> (T input, ReductionOp reduction_op)</td></tr>
<tr class="memdesc:a1aa0a6e1a2c3feabbc20c27864df15c1"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes a warp-wide reduction in each active warp using the specified binary reduction functor. The output is valid in warp <em>lane</em><sub>0</sub>.  <a href="#a1aa0a6e1a2c3feabbc20c27864df15c1">More...</a><br/></td></tr>
<tr class="separator:a1aa0a6e1a2c3feabbc20c27864df15c1"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a176c3c913e3e077aab691a186f161d8e"><td class="memTemplParams" colspan="2">template&lt;typename ReductionOp &gt; </td></tr>
<tr class="memitem:a176c3c913e3e077aab691a186f161d8e"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ T&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_reduce.html#a176c3c913e3e077aab691a186f161d8e">Reduce</a> (T input, ReductionOp reduction_op, int valid_items)</td></tr>
<tr class="memdesc:a176c3c913e3e077aab691a186f161d8e"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes a partially-full warp-wide reduction in each active warp using the specified binary reduction functor. The output is valid in warp <em>lane</em><sub>0</sub>.  <a href="#a176c3c913e3e077aab691a186f161d8e">More...</a><br/></td></tr>
<tr class="separator:a176c3c913e3e077aab691a186f161d8e"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a85d8612201876163a74a807c585b0af1"><td class="memTemplParams" colspan="2">template&lt;typename ReductionOp , typename Flag &gt; </td></tr>
<tr class="memitem:a85d8612201876163a74a807c585b0af1"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ T&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_reduce.html#a85d8612201876163a74a807c585b0af1">HeadSegmentedReduce</a> (T input, Flag head_flag, ReductionOp reduction_op)</td></tr>
<tr class="memdesc:a85d8612201876163a74a807c585b0af1"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes a segmented reduction in each active warp where segments are defined by head-flags. The reduction of each segment is returned to the first lane in that segment (which always includes <em>lane</em><sub>0</sub>).  <a href="#a85d8612201876163a74a807c585b0af1">More...</a><br/></td></tr>
<tr class="separator:a85d8612201876163a74a807c585b0af1"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a074ddcaf999d4e97376b145086e78ddb"><td class="memTemplParams" colspan="2">template&lt;typename ReductionOp , typename Flag &gt; </td></tr>
<tr class="memitem:a074ddcaf999d4e97376b145086e78ddb"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ T&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_reduce.html#a074ddcaf999d4e97376b145086e78ddb">TailSegmentedReduce</a> (T input, Flag tail_flag, ReductionOp reduction_op)</td></tr>
<tr class="memdesc:a074ddcaf999d4e97376b145086e78ddb"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes a segmented reduction in each active warp where segments are defined by tail-flags. The reduction of each segment is returned to the first lane in that segment (which always includes <em>lane</em><sub>0</sub>).  <a href="#a074ddcaf999d4e97376b145086e78ddb">More...</a><br/></td></tr>
<tr class="separator:a074ddcaf999d4e97376b145086e78ddb"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table>
<h2 class="groupheader">Constructor &amp; Destructor Documentation</h2>
<a class="anchor" id="a7e026682e53d5af519912ab9a9215e3b"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ <a class="el" href="classcub_1_1_warp_reduce.html">cub::WarpReduce</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::<a class="el" href="classcub_1_1_warp_reduce.html">WarpReduce</a> </td>
          <td>(</td>
          <td class="paramtype"><a class="el" href="structcub_1_1_warp_reduce_1_1_temp_storage.html">TempStorage</a> &amp;&#160;</td>
          <td class="paramname"><em>temp_storage</em>)</td><td></td>
          <td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Collective constructor for 1D thread blocks using the specified memory allocation as temporary storage. Logical warp and lane identifiers are constructed from <code>threadIdx.x</code>. </p>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">temp_storage</td><td>Reference to memory allocation having layout type <a class="el" href="structcub_1_1_warp_reduce_1_1_temp_storage.html" title="The operations exposed by WarpReduce require a temporary memory allocation of this nested type for th...">TempStorage</a> </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__reduce_8cuh_source.html#l00229">229</a> of file <a class="el" href="warp__reduce_8cuh_source.html">warp_reduce.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="aef629abf014a0f1f9646d3b9123fcaa6"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ <a class="el" href="classcub_1_1_warp_reduce.html">cub::WarpReduce</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::<a class="el" href="classcub_1_1_warp_reduce.html">WarpReduce</a> </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>warp_id</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>lane_id</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Collective constructor using a private static allocation of shared memory as temporary storage. Threads are identified using the given warp and lane identifiers. </p>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">warp_id</td><td>A suitable warp membership identifier </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">lane_id</td><td>A lane identifier within the warp </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__reduce_8cuh_source.html#l00245">245</a> of file <a class="el" href="warp__reduce_8cuh_source.html">warp_reduce.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="aac44f5e8003911f074156366969947e4"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ <a class="el" href="classcub_1_1_warp_reduce.html">cub::WarpReduce</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::<a class="el" href="classcub_1_1_warp_reduce.html">WarpReduce</a> </td>
          <td>(</td>
          <td class="paramtype"><a class="el" href="structcub_1_1_warp_reduce_1_1_temp_storage.html">TempStorage</a> &amp;&#160;</td>
          <td class="paramname"><em>temp_storage</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>warp_id</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>lane_id</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Collective constructor using the specified memory allocation as temporary storage. Threads are identified using the given warp and lane identifiers. </p>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">temp_storage</td><td>Reference to memory allocation having layout type <a class="el" href="structcub_1_1_warp_reduce_1_1_temp_storage.html" title="The operations exposed by WarpReduce require a temporary memory allocation of this nested type for th...">TempStorage</a> </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">warp_id</td><td>A suitable warp membership identifier </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">lane_id</td><td>A lane identifier within the warp </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__reduce_8cuh_source.html#l00258">258</a> of file <a class="el" href="warp__reduce_8cuh_source.html">warp_reduce.cuh</a>.</p>

</div>
</div>
<h2 class="groupheader">Member Function Documentation</h2>
<a class="anchor" id="a85d8612201876163a74a807c585b0af1"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<div class="memtemplate">
template&lt;typename ReductionOp , typename Flag &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ T <a class="el" href="classcub_1_1_warp_reduce.html">cub::WarpReduce</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::HeadSegmentedReduce </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">Flag&#160;</td>
          <td class="paramname"><em>head_flag</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">ReductionOp&#160;</td>
          <td class="paramname"><em>reduction_op</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes a segmented reduction in each active warp where segments are defined by head-flags. The reduction of each segment is returned to the first lane in that segment (which always includes <em>lane</em><sub>0</sub>). </p>
<p>Supports non-commutative reduction operators</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates a head-segmented warp max reduction within a block of 32 threads (one warp). </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpReduce for a single warp on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_reduce.html" title="The WarpReduce class provides collective methods for computing a parallel reduction of items partitio...">cub::WarpReduce&lt;int, 1&gt;</a> <a class="code" href="classcub_1_1_warp_reduce.html#aed6c8ead3a4c2fa56ae1d5851fb36848" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpReduce</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpReduce</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpReduce::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain one input item and flag per thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data = ...</div>
<div class="line">    <span class="keywordtype">int</span> head_flag = ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Return the warp-wide reductions to each lane0</span></div>
<div class="line">    <span class="keywordtype">int</span> aggregate = <a class="code" href="classcub_1_1_warp_reduce.html#aed6c8ead3a4c2fa56ae1d5851fb36848" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpReduce</a>(temp_storage).HeadSegmentedReduce(</div>
<div class="line">        thread_data, head_flag, <a class="code" href="structcub_1_1_max.html" title="Default max functor. ">cub::Max</a>());</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> and <code>head_flag</code> across the block of threads is <code>0, 1, 2, 3, ..., 31</code> and is <code>1, 0, 0, 0, 1, 0, 0, 0, ..., 1, 0, 0, 0</code>, respectively. The corresponding output <code>aggregate</code> in threads 0, 4, 8, etc. will be <code>3</code>, <code>7</code>, <code>11</code>, etc. (and is undefined in other threads).</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ReductionOp</td><td><b>[inferred]</b> Binary reduction operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">head_flag</td><td>Head flag denoting whether or not <code>input</code> is the start of a new segment </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">reduction_op</td><td>Reduction operator </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__reduce_8cuh_source.html#l00611">611</a> of file <a class="el" href="warp__reduce_8cuh_source.html">warp_reduce.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="abc0cd1853c09b27fc45f564596163be4"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<div class="memtemplate">
template&lt;typename Flag &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ T <a class="el" href="classcub_1_1_warp_reduce.html">cub::WarpReduce</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::HeadSegmentedSum </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">Flag&#160;</td>
          <td class="paramname"><em>head_flag</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes a segmented sum in each active warp where segments are defined by head-flags. The sum of each segment is returned to the first lane in that segment (which always includes <em>lane</em><sub>0</sub>). </p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates a head-segmented warp sum reduction within a block of 32 threads (one warp). </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpReduce for a single warp on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_reduce.html" title="The WarpReduce class provides collective methods for computing a parallel reduction of items partitio...">cub::WarpReduce&lt;int, 1&gt;</a> <a class="code" href="classcub_1_1_warp_reduce.html#aed6c8ead3a4c2fa56ae1d5851fb36848" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpReduce</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpReduce</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpReduce::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain one input item and flag per thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data = ...</div>
<div class="line">    <span class="keywordtype">int</span> head_flag = ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Return the warp-wide sums to each lane0</span></div>
<div class="line">    <span class="keywordtype">int</span> aggregate = <a class="code" href="classcub_1_1_warp_reduce.html#aed6c8ead3a4c2fa56ae1d5851fb36848" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpReduce</a>(temp_storage).HeadSegmentedSum(</div>
<div class="line">        thread_data, head_flag);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> and <code>head_flag</code> across the block of threads is <code>0, 1, 2, 3, ..., 31</code> and is <code>1, 0, 0, 0, 1, 0, 0, 0, ..., 1, 0, 0, 0</code>, respectively. The corresponding output <code>aggregate</code> in threads 0, 4, 8, etc. will be <code>6</code>, <code>22</code>, <code>38</code>, etc. (and is undefined in other threads).</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ReductionOp</td><td><b>[inferred]</b> Binary reduction operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">head_flag</td><td>Head flag denoting whether or not <code>input</code> is the start of a new segment </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__reduce_8cuh_source.html#l00407">407</a> of file <a class="el" href="warp__reduce_8cuh_source.html">warp_reduce.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a1aa0a6e1a2c3feabbc20c27864df15c1"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<div class="memtemplate">
template&lt;typename ReductionOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ T <a class="el" href="classcub_1_1_warp_reduce.html">cub::WarpReduce</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::Reduce </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">ReductionOp&#160;</td>
          <td class="paramname"><em>reduction_op</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes a warp-wide reduction in each active warp using the specified binary reduction functor. The output is valid in warp <em>lane</em><sub>0</sub>. </p>
<p>Supports non-commutative reduction operators</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates four concurrent warp max reductions within a block of 128 threads (one per each of the 32-thread warps). </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpReduce for 4 warps on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_reduce.html" title="The WarpReduce class provides collective methods for computing a parallel reduction of items partitio...">cub::WarpReduce&lt;int, 4&gt;</a> <a class="code" href="classcub_1_1_warp_reduce.html#aed6c8ead3a4c2fa56ae1d5851fb36848" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpReduce</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpReduce</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpReduce::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain one input item per thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data = ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Return the warp-wide reductions to each lane0</span></div>
<div class="line">    <span class="keywordtype">int</span> aggregate = <a class="code" href="classcub_1_1_warp_reduce.html#aed6c8ead3a4c2fa56ae1d5851fb36848" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpReduce</a>(temp_storage).Reduce(</div>
<div class="line">        thread_data, <a class="code" href="structcub_1_1_max.html" title="Default max functor. ">cub::Max</a>());</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>0, 1, 2, 3, ..., 127</code>. The corresponding output <code>aggregate</code> in threads 0, 32, 64, and 96 will <code>31</code>, <code>63</code>, <code>95</code>, and <code>127</code>, respectively (and is undefined in other threads).</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ReductionOp</td><td><b>[inferred]</b> Binary reduction operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">reduction_op</td><td>Binary reduction operator </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__reduce_8cuh_source.html#l00505">505</a> of file <a class="el" href="warp__reduce_8cuh_source.html">warp_reduce.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a176c3c913e3e077aab691a186f161d8e"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<div class="memtemplate">
template&lt;typename ReductionOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ T <a class="el" href="classcub_1_1_warp_reduce.html">cub::WarpReduce</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::Reduce </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">ReductionOp&#160;</td>
          <td class="paramname"><em>reduction_op</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>valid_items</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes a partially-full warp-wide reduction in each active warp using the specified binary reduction functor. The output is valid in warp <em>lane</em><sub>0</sub>. </p>
<p>All threads in each logical warp must agree on the same value for <code>valid_items</code>. Otherwise the result is undefined.</p>
<p>Supports non-commutative reduction operators</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates a max reduction within a single, partially-full block of 32 threads (one warp). </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(<span class="keywordtype">int</span> *d_data, <span class="keywordtype">int</span> valid_items)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpReduce for a single warp on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_reduce.html" title="The WarpReduce class provides collective methods for computing a parallel reduction of items partitio...">cub::WarpReduce&lt;int, 1&gt;</a> <a class="code" href="classcub_1_1_warp_reduce.html#aed6c8ead3a4c2fa56ae1d5851fb36848" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpReduce</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpReduce</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpReduce::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain one input item per thread if in range</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data;</div>
<div class="line">    <span class="keywordflow">if</span> (threadIdx.x &lt; valid_items)</div>
<div class="line">        thread_data = d_data[threadIdx.x];</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Return the warp-wide reductions to each lane0</span></div>
<div class="line">    <span class="keywordtype">int</span> aggregate = <a class="code" href="classcub_1_1_warp_reduce.html#aed6c8ead3a4c2fa56ae1d5851fb36848" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpReduce</a>(temp_storage).Reduce(</div>
<div class="line">        thread_data, <a class="code" href="structcub_1_1_max.html" title="Default max functor. ">cub::Max</a>(), valid_items);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the input <code>d_data</code> is <code>0, 1, 2, 3, 4, ...</code> and <code>valid_items</code> is <code>4</code>. The corresponding output <code>aggregate</code> in thread0 is <code>3</code> (and is undefined in other threads).</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ReductionOp</td><td><b>[inferred]</b> Binary reduction operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">reduction_op</td><td>Binary reduction operator </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">valid_items</td><td>Total number of valid items in the calling thread's logical warp (may be less than <code>LOGICAL_WARP_THREADS</code>) </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__reduce_8cuh_source.html#l00553">553</a> of file <a class="el" href="warp__reduce_8cuh_source.html">warp_reduce.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a3327a2b8c9bdce58f7af1e2485489eed"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ T <a class="el" href="classcub_1_1_warp_reduce.html">cub::WarpReduce</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::<a class="el" href="structcub_1_1_sum.html">Sum</a> </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>)</td><td></td>
          <td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes a warp-wide sum in each active warp. The output is valid in warp <em>lane</em><sub>0</sub>. </p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates four concurrent warp sum reductions within a block of 128 threads (one per each of the 32-thread warps). </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpReduce for 4 warps on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_reduce.html" title="The WarpReduce class provides collective methods for computing a parallel reduction of items partitio...">cub::WarpReduce&lt;int, 4&gt;</a> <a class="code" href="classcub_1_1_warp_reduce.html#aed6c8ead3a4c2fa56ae1d5851fb36848" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpReduce</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpReduce</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpReduce::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain one input item per thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data = ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Return the warp-wide sums to each lane0</span></div>
<div class="line">    <span class="keywordtype">int</span> aggregate = <a class="code" href="classcub_1_1_warp_reduce.html#aed6c8ead3a4c2fa56ae1d5851fb36848" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpReduce</a>(temp_storage).Sum(thread_data);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>0, 1, 2, 3, ..., 127</code>. The corresponding output <code>aggregate</code> in threads 0, 32, 64, and 96 will <code>496</code>, <code>1520</code>, <code>2544</code>, and <code>3568</code>, respectively (and is undefined in other threads). </dd></dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__reduce_8cuh_source.html#l00309">309</a> of file <a class="el" href="warp__reduce_8cuh_source.html">warp_reduce.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a2ccc6eb62a1e6caf145eb94a8cd58f0f"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ T <a class="el" href="classcub_1_1_warp_reduce.html">cub::WarpReduce</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::<a class="el" href="structcub_1_1_sum.html">Sum</a> </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>valid_items</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes a partially-full warp-wide sum in each active warp. The output is valid in warp <em>lane</em><sub>0</sub>. </p>
<p>All threads in each logical warp must agree on the same value for <code>valid_items</code>. Otherwise the result is undefined.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates a sum reduction within a single, partially-full block of 32 threads (one warp). </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(<span class="keywordtype">int</span> *d_data, <span class="keywordtype">int</span> valid_items)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpReduce for a single warp on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_reduce.html" title="The WarpReduce class provides collective methods for computing a parallel reduction of items partitio...">cub::WarpReduce&lt;int, 1&gt;</a> <a class="code" href="classcub_1_1_warp_reduce.html#aed6c8ead3a4c2fa56ae1d5851fb36848" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpReduce</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpReduce</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpReduce::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain one input item per thread if in range</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data;</div>
<div class="line">    <span class="keywordflow">if</span> (threadIdx.x &lt; valid_items)</div>
<div class="line">        thread_data = d_data[threadIdx.x];</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Return the warp-wide sums to each lane0</span></div>
<div class="line">    <span class="keywordtype">int</span> aggregate = <a class="code" href="classcub_1_1_warp_reduce.html#aed6c8ead3a4c2fa56ae1d5851fb36848" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpReduce</a>(temp_storage).Sum(</div>
<div class="line">        thread_data, valid_items);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the input <code>d_data</code> is <code>0, 1, 2, 3, 4, ...</code> and <code>valid_items</code> is <code>4</code>. The corresponding output <code>aggregate</code> in thread0 is <code>6</code> (and is undefined in other threads). </dd></dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">valid_items</td><td>Total number of valid items in the calling thread's logical warp (may be less than <code>LOGICAL_WARP_THREADS</code>) </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__reduce_8cuh_source.html#l00352">352</a> of file <a class="el" href="warp__reduce_8cuh_source.html">warp_reduce.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a074ddcaf999d4e97376b145086e78ddb"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<div class="memtemplate">
template&lt;typename ReductionOp , typename Flag &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ T <a class="el" href="classcub_1_1_warp_reduce.html">cub::WarpReduce</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::TailSegmentedReduce </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">Flag&#160;</td>
          <td class="paramname"><em>tail_flag</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">ReductionOp&#160;</td>
          <td class="paramname"><em>reduction_op</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes a segmented reduction in each active warp where segments are defined by tail-flags. The reduction of each segment is returned to the first lane in that segment (which always includes <em>lane</em><sub>0</sub>). </p>
<p>Supports non-commutative reduction operators</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates a tail-segmented warp max reduction within a block of 32 threads (one warp). </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpReduce for a single warp on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_reduce.html" title="The WarpReduce class provides collective methods for computing a parallel reduction of items partitio...">cub::WarpReduce&lt;int, 1&gt;</a> <a class="code" href="classcub_1_1_warp_reduce.html#aed6c8ead3a4c2fa56ae1d5851fb36848" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpReduce</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpReduce</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpReduce::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain one input item and flag per thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data = ...</div>
<div class="line">    <span class="keywordtype">int</span> tail_flag = ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Return the warp-wide reductions to each lane0</span></div>
<div class="line">    <span class="keywordtype">int</span> aggregate = <a class="code" href="classcub_1_1_warp_reduce.html#aed6c8ead3a4c2fa56ae1d5851fb36848" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpReduce</a>(temp_storage).TailSegmentedReduce(</div>
<div class="line">        thread_data, tail_flag, <a class="code" href="structcub_1_1_max.html" title="Default max functor. ">cub::Max</a>());</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> and <code>tail_flag</code> across the block of threads is <code>0, 1, 2, 3, ..., 31</code> and is <code>0, 0, 0, 1, 0, 0, 0, 1, ..., 0, 0, 0, 1</code>, respectively. The corresponding output <code>aggregate</code> in threads 0, 4, 8, etc. will be <code>3</code>, <code>7</code>, <code>11</code>, etc. (and is undefined in other threads).</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ReductionOp</td><td><b>[inferred]</b> Binary reduction operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">tail_flag</td><td>Tail flag denoting whether or not <code>input</code> is the end of the current segment </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">reduction_op</td><td>Reduction operator </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__reduce_8cuh_source.html#l00661">661</a> of file <a class="el" href="warp__reduce_8cuh_source.html">warp_reduce.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a5b0134131d03909e43d24d6b0b50beb4"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<div class="memtemplate">
template&lt;typename Flag &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ T <a class="el" href="classcub_1_1_warp_reduce.html">cub::WarpReduce</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::TailSegmentedSum </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">Flag&#160;</td>
          <td class="paramname"><em>tail_flag</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes a segmented sum in each active warp where segments are defined by tail-flags. The sum of each segment is returned to the first lane in that segment (which always includes <em>lane</em><sub>0</sub>). </p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates a tail-segmented warp sum reduction within a block of 32 threads (one warp). </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpReduce for a single warp on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_reduce.html" title="The WarpReduce class provides collective methods for computing a parallel reduction of items partitio...">cub::WarpReduce&lt;int, 1&gt;</a> <a class="code" href="classcub_1_1_warp_reduce.html#aed6c8ead3a4c2fa56ae1d5851fb36848" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpReduce</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpReduce</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpReduce::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain one input item and flag per thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data = ...</div>
<div class="line">    <span class="keywordtype">int</span> tail_flag = ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Return the warp-wide sums to each lane0</span></div>
<div class="line">    <span class="keywordtype">int</span> aggregate = <a class="code" href="classcub_1_1_warp_reduce.html#aed6c8ead3a4c2fa56ae1d5851fb36848" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpReduce</a>(temp_storage).TailSegmentedSum(</div>
<div class="line">        thread_data, tail_flag);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> and <code>tail_flag</code> across the block of threads is <code>0, 1, 2, 3, ..., 31</code> and is <code>0, 0, 0, 1, 0, 0, 0, 1, ..., 0, 0, 0, 1</code>, respectively. The corresponding output <code>aggregate</code> in threads 0, 4, 8, etc. will be <code>6</code>, <code>22</code>, <code>38</code>, etc. (and is undefined in other threads).</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ReductionOp</td><td><b>[inferred]</b> Binary reduction operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">tail_flag</td><td>Head flag denoting whether or not <code>input</code> is the start of a new segment </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__reduce_8cuh_source.html#l00453">453</a> of file <a class="el" href="warp__reduce_8cuh_source.html">warp_reduce.cuh</a>.</p>

</div>
</div>
<hr/>The documentation for this class was generated from the following file:<ul>
<li><a class="el" href="warp__reduce_8cuh_source.html">warp_reduce.cuh</a></li>
</ul>
</div><!-- contents -->
<!-- HTML footer for doxygen 1.8.3.1-->
<!-- start footer part -->
<hr class="footer"/><address class="footer"><small>
Generated on Fri Aug 23 2013 17:31:13 for CUB by &#160;<a href="http://www.doxygen.org/index.html">
<img class="footer" src="doxygen.png" alt="doxygen"/>
</a> 1.8.4
<br>
&copy; 2013 NVIDIA Corporation
</small></address>
</body>
</html>
